#ifndef _RKNN_DEMO_DEEPLABV3_KERNEL_H_
#define _RKNN_DEMO_DEEPLABV3_KERNEL_H_

//"#define SRC_STRIDE 1365 //65*21\n" 
const char* CL_kernel_string_src = 

"#define INC(x,l) min(x+1,l-1)\n"
"__kernel void UpsampleSoftmax(__global const float* src_buf, __global uchar* index_buf, "
"                              const int srcHeight, const int srcWidth, "
"                              const int dstHeight, const int dstWidth, const int dstChannel,"
"                              const float scale_h_inv, const float scale_w_inv, const int SRC_STRIDE) {\n"
"  int dx = get_global_id(0);//dst width \n"
"  int dy = get_global_id(1);//dst height \n"
"  float sx = ((dx) * scale_w_inv), sy = ((dy) * scale_h_inv); //resize\n"
"  int x = floor(sx), y = floor(sy); \n"
"  float u = sx - x, v = sy - y; \n"
"  if ( x<0 ) x=0,u=0; \n"
"  if ( x>=srcWidth ) x=srcWidth-1,u=0; \n"
"  if ( y<0 ) y=0,v=0; \n"
"  if ( y>=srcHeight ) y=srcHeight-1,v=0; \n"
"  int y_ = INC(y, srcHeight); \n"
"  int x_ = INC(x, srcWidth); \n"
"  float u1 = 1.f - u; \n"
"  float v1 = 1.f - v; \n"
"  float w0 = u1 * v1; \n"
"  float w1 = u * v1; \n"
"  float w2 = u1 * v; \n"
"  float w3 = u * v; \n"
"  int index0 = mad24(y, SRC_STRIDE, mul24(x, dstChannel)); \n"
"  int index1 = mad24(y, SRC_STRIDE, mul24(x_, dstChannel)); \n"
"  int index2 = mad24(y_, SRC_STRIDE, mul24(x, dstChannel)); \n"
"  int index3 = mad24(y_, SRC_STRIDE, mul24(x_, dstChannel)); \n"
"  float data0 = src_buf[index0]; \n"
"  float data1 = src_buf[index1]; \n"
"  float data2 = src_buf[index2]; \n"
"  float data3 = src_buf[index3]; \n"
"  float max_val = fma(w0, data0, fma(w1,data1, fma(w2, data2, fma(w3,data3,0.0f)))); \n"
"  int4 index = {index0, index1, index2, index3}; \n"
"  index+=1; \n"
"  int group_num = 5; \n"
"  int max_group_id = -1; \n"
"  int max_v_id = 0; \n"
"  for(int i = 0; i< group_num;++i) { \n"
"   float4 data0_v4 = vload4(0, src_buf+index.x); \n"
"   float4 data1_v4 = vload4(0, src_buf+index.y); \n"
"   float4 data2_v4 = vload4(0, src_buf+index.z); \n"
"   float4 data3_v4 = vload4(0, src_buf+index.w); \n"
"   float4 u_val = fma(w0, data0_v4, fma(w1,data1_v4, fma(w2, data2_v4, fma(w3,data3_v4,0.0f)))); \n"
"   float max_xy_v = fmax(u_val.x, u_val.y); \n"
"   float max_zw_v = fmax(u_val.z, u_val.w); \n"
"   float temp_max_val = fmax(max_xy_v, max_zw_v); \n"
"   if(temp_max_val > max_val) { \n"
"     max_val = temp_max_val; \n"
"    int max_xy_id = select(0,1,isless(u_val.x, u_val.y)); \n"
"     int max_zw_id = select(2,3,isless(u_val.z, u_val.w)); \n"
"     max_v_id = select(max_xy_id, max_zw_id, isless(max_xy_v, max_zw_v)); \n"
"     max_group_id = i; }\n"
"    index+=4; } \n"
"  __global uchar *dst_ptr = index_buf + mad24(dy,dstWidth,dx); \n"
"  dst_ptr[0] = select(max_group_id*4+max_v_id, 0, max_group_id == -1);\n"
"}";

#endif //_RKNN_DEMO_DEEPLABV3_H_